In `@testing/python/language/test_tilelang_language_reduce.py`:
- Around line 99-105: The ref_fn uses torch.uint32 in a static dtype list which
raises AttributeError on PyTorch <2.3.0; update ref_fn to conditionally include
torch.uint32 only when hasattr(torch, "uint32") (or use getattr with a fallback)
so the dtype check is built at runtime, mirroring the existing pattern used for
version-dependent dtypes like float8_e4m3fn; locate and modify the dtype
membership test in ref_fn to construct the list/set conditionally and then
perform the same res.to(A.dtype) conversion for the supported integer dtypes
(torch.uint32, torch.int32, torch.int64).

In `@src/tl_templates/cuda/reduce.h`:
- Around line 377-381: The fast-path incorrectly casts 64-bit integrals to
int32_t and uses run_reduce_sync, truncating results; update the compile-time
branch in the template so the run_reduce_sync path is only taken when T is
exactly int32_t or uint32_t (remove the branch that treats other
std::is_integral_v<T> types by casting), and let 64-bit types fall through to
the shfl_xor_sync-based fallback that correctly handles 64-bit operands; adjust
the condition around run_reduce_sync (the constexpr if that currently checks
std::is_same_v<T, int32_t> || std::is_same_v<T, uint32_t> and the following
std::is_integral_v<T> branch) accordingly so no cast-to-int32_t is performed for
64-bit integrals.
- Around line 344-358: The sm100 fast-path incorrectly casts floats and calls
__reduce_max_sync/__reduce_min_sync which only accept 32-bit integer operands;
modify the branch in reduce.h so the __reduce_*_sync calls are only compiled for
32-bit integer types (e.g., when std::is_same_v<T,int> or
std::is_same_v<T,unsigned int> / std::is_integral and sizeof(T)==4), and remove
the float/half/bfloat casting there; for all other types (float, half_t,
bfloat16_t, etc.) fall through to the existing non-sm100 handling. Ensure you
static_cast the value to int/unsigned before passing to
__reduce_max_sync/__reduce_min_sync and keep the ReduceOp/MaxOp/MinOp constexpr
checks to select the right intrinsic.

In `@src/tl_templates/cuda/reduce.h`:
- Around line 344-366: The SM100 asm fast-path currently casts any T to float
and uses redux.sync.*.f32, which corrupts large integer ordering; change the
compile-time guard inside the SM100 block to only take this fast-path for
floating types (float, half_t, bfloat16_t) — e.g. wrap the asm branches with an
if constexpr that checks std::is_floating_point_v<T> || std::is_same_v<T,
half_t> || std::is_same_v<T, bfloat16_t> so integral types fall through to the
SM80+ __reduce_max_sync/__reduce_min_sync implementations; keep the existing
value_cast and asm/result handling for the floating case and do not alter the
fallback path for non-floating T.

In `@src/tl_templates/cuda/reduce.h`:
- Around line 344-365: The fast-path in the CUDA reduce code leaves value_cast
as 0.0f for T=float (and mishandles double), so the inline asm reduces zeros;
update the branches in the block inside reduce.h (the section that sets
value_cast and emits redux.sync) to handle float explicitly by assigning
value_cast = value when std::is_same_v<T, float>, keep the existing assignments
for half_t and bfloat16_t, and ensure the fast-path is enabled only for the
three safe types (float, half_t, bfloat16_t) — do NOT include double — before
emitting the MaxOp/MinOp asm so the asm uses the real value_cast.
- Around line 368-382: The template lambda run_reduce_sync uses C++20-only
syntax and must be replaced for C++17: implement a regular templated helper
function (or a small function object) named e.g. run_reduce_sync_impl templated
on typename T_cast that contains the same if constexpr branches for ReduceOp vs
SumOp/MaxOp/MinOp/BitAndOp/BitOrOp/BitXorOp and calls the corresponding
__reduce_*_sync(mask, val), then call that helper from the existing site instead
of the lambda; ensure the helper has the same capture visibility (pass mask in
as a parameter or make it accessible) and is declared in the same scope so
ReduceOp and the __reduce_*_sync calls resolve correctly.

In `@testing/python/language/test_tilelang_language_reduce.py`:
- Around line 46-61: The test's dispatch on op (variable op) lacks a final else
branch so unsupported op values leave B_local uninitialized and produce silent
garbage; update the branch in the block that calls T.reduce_sum / T.reduce_max /
T.reduce_min / T.reduce_abssum / T.reduce_absmax / T.reduce_bitand /
T.reduce_bitor / T.reduce_bitxor to include a final else that raises an explicit
exception (e.g., raise ValueError or AssertionError) including the invalid op
string so the test fails loudly and identifies the unsupported operation.
- Around line 98-111: ref_fn currently lacks branches for the bitwise ops
(bitand, bitor, bitxor) so res can be unassigned; add branches inside ref_fn to
handle "bitand", "bitor", and "bitxor" by performing a column-wise reduction
over dim=1 (looping over columns and applying the appropriate
torch.bitwise_and/bitwise_or/bitwise_xor) and for "bitand" initialize the
accumulator with an all-ones tensor of the same dtype/shape as a row (use
~torch.zeros_like(row) or equivalent) so types like torch.uint32/int32/int64
behave correctly; keep the existing dtype-preservation logic (the check using
A.dtype in [...]) and return the reduced tensor cast back when needed.

In `@src/op/reduce.cc`:
- Around line 392-408: The hierarchical-workspace sizing is wrong when
reducing_threads < all_threads: is_hierarchical currently makes workspace_size =
reducing_threads / 32, but threads use global_warp_id = (threadIdx.x -
thread_offset) / 32 which can index up to (all_threads-1)/32, causing
out-of-bounds writes in hierarchical_reduce; fix by sizing the workspace to
cover all warps in the block when is_hierarchical (use the block thread extent
from T.thread_bounds->extent, e.g. workspace_size =
static_cast<int>(*as_const_int(T.thread_bounds->extent) + 31) / 32 or equivalent
integer division rounding up), and keep using T.AddWorkspace and
thread_reduce_args as before (symbols to edit: is_hierarchical,
reducing_threads, T.thread_bounds->extent, T.AddWorkspace, hierarchical_reduce,
global_warp_id, thread_offset).

In `@src/tl_templates/cuda/reduce.h`:
- Around line 65-78: The comment notes that hierarchical_reduce now uses barrier
phases 1, 2 and 3, so barrier IDs 1–3 are reserved for internal use; update any
user-facing barrier documentation and comments that currently state "user
barriers start from 3" to instead state that user-defined barrier IDs must start
from 4, and add a short note near the barrier policy declarations
(SyncThreadsBarrier and NamedBarrier) and in the hierarchical_reduce
documentation mentioning that hierarchical_reduce occupies phases 1–3 to avoid
future sync conflicts.
- Around line 140-164: hierarchical_reduce can write past red_buf when the
caller uses fewer "reducing_threads" than the full block (global_warp_id may
exceed reducing_threads/32); fix by adding bounds checks or by having the caller
allocate red_buf sized to total global warps: ensure before writing
red_buf[global_warp_id] and before returning/reading red_buf[group_id *
num_warps_per_group] that global_warp_id and group_id*... are within the
allocated workspace (or clamp/skips writes for warps >= reducing_warps), and
verify warp_inter_reduce is only invoked for groups that have valid entries;
update the caller that allocates red_buf (and any code that computes
reducing_threads) to size it to ceil(reducing_threads/32) or to total global
warps consistently.